#include<stdio.h>
#include <cutil_inline.h>
#include<assert.h>
#include<stdlib.h>
#include<time.h>

#define ARRAY_SIZE 8388608
#define THREADS_PER_BLOCK 256

#define NSTREAMS 4
#define NBLOCKS ARRAY_SIZE / (THREADS_PER_BLOCK * NSTREAMS)
#define ITER 10


__global__ void
execute(int* const dData)
{
    const int thread = blockIdx.x *  blockDim.x + threadIdx.x;
    dData[thread] += 10;
}


__global__ void
dummy()
{
    __shared__ int  temp;
    temp = 3;
    temp = temp + 1;
}

int
main()
{
    const unsigned size = sizeof(int) * ARRAY_SIZE;
    int* data = 0;
    cutilSafeCall( cudaMallocHost((void**)&data, size));

    int* dData = 0;
    cutilSafeCall(cudaMalloc((void**)&dData, size));

    dim3 threads(THREADS_PER_BLOCK);
    dim3 blocks(NBLOCKS);
    dummy <<< blocks, threads >>> ();
    

     for (unsigned i = 0; i < ARRAY_SIZE; ++i) {
          data[i] = 0;
     }

     cudaStream_t stream[NSTREAMS];
     for(unsigned  i = 0; i < NSTREAMS; ++i) {
         cutilSafeCall(cudaStreamCreate(&stream[i]));
     }
    
     const int chunkSize = size / NSTREAMS; 
     cudaEvent_t startEvent, stopEvent;
     cutilSafeCall( cudaEventCreate(&startEvent));
     cutilSafeCall( cudaEventCreate(&stopEvent) );



     cudaEventRecord(startEvent, 0);
     for (unsigned k = 0; k < ITER; ++k) {
		 for (unsigned i = 0; i < NSTREAMS; ++i) {

			 const int offset = i * ARRAY_SIZE / NSTREAMS;
			 cudaMemcpyAsync(dData + offset, data + offset, chunkSize,
					 cudaMemcpyHostToDevice, stream[i]);
		 }

		 for (unsigned i = 0; i < NSTREAMS; ++i) {
			 const int offset = i * ARRAY_SIZE / NSTREAMS;
			 execute <<< blocks, threads, 0, stream[i]  >>> (dData + offset );
		 }

		 for (unsigned i = 0; i < NSTREAMS; ++i) {
			 const int offset = i * ARRAY_SIZE / NSTREAMS;
			 cudaMemcpyAsync(data + offset, dData + offset, chunkSize,
					 cudaMemcpyDeviceToHost, stream[i]);
		 }
     }

     cudaEventRecord(stopEvent, 0);
     cudaEventSynchronize(stopEvent);
     float elapsedTime = 0;
     cutilSafeCall( cudaEventElapsedTime(&elapsedTime, 
                    startEvent, stopEvent));
     cudaEventDestroy(startEvent);
     cudaEventDestroy(stopEvent);


     printf("Processing time for asynch %f\n", elapsedTime); 
 


     cutilSafeCall(cudaFree(dData));
      

     for (unsigned i = 0; i < ARRAY_SIZE; ++i) {
         assert(data[i] == ITER * 10);
     }

     for(unsigned  i = 0; i < NSTREAMS; ++i) {
         cutilSafeCall(cudaStreamDestroy(stream[i]));
     }

     cudaFreeHost(data);
}

